#include <math.h> // function to add the elements of two arrays
#include <iostream>
/*
https://developer.nvidia.com/blog/even-easier-introduction-cuda/
nomentation:
    kernels (functions)
    device code
    host code
    Unified Memory 
    launch (the kernel)

    race condition 
    threads in a thread block
    nvprof --> nsys profile ./add_cuda --> 
    generate a summary of kernel execution times and other statistics:  nsys stats ./report2.nsys-rep
    cuda_gpu_kern_sum

    Streaming Multiprocessors, or SMs.
    thread blocks.
    grid-stride loop.
*/

// CUDA Kernel function to add the elements of two arrays on the GPU
// 每个线程执行一次计算，而不是将计算分散到并行线程上。
__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

int main(void)
{
    int    N = 1 << 20; // 表示将二进制数 1 左移 20 位。

    // Allocate Unified Memory -- accessible from CPU or GPU
    float *x, *y;
    cudaMallocManaged(&x, N*sizeof(float));
    cudaMallocManaged(&y, N*sizeof(float));
    
    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++)
    {
        x[i] = 1.0f;
        y[i] = 2.0f;
    } 

    //  最佳划分blocks和每个blocks的线程数的实践示例：
    int threadsPerBlock = 256; // 是 32 的倍数，且 << 1024
    int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
    add<<<numBlocks, threadsPerBlock>>>(N, x, y);
    
    // CUDA kernel launches don’t block the calling CPU thread, so call cudaDeviceSynchronize()
    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();

    // Check for errors (all values should be 3.0f) 
    float maxError = 0.0f; 
    for (int i = 0; i < N; i++) 
        maxError = fmax(maxError, fabs(y[i]-3.0f));
    
    std::cout << "Max error: " << maxError << std::endl; 
    
    // Free memory
    cudaFree(x);
    cudaFree(y);
    
    return 0;
}